kernel void op_3 (
    device float* out [[ buffer(0) ]],
    const device float* in0 [[ buffer(1) ]],
    const device float* in1 [[ buffer(2) ]],
    uint gid [[ thread_position_in_grid ]])
{
    constant long shape[] = {2,3};
    constant long strides_in0[] = {3,1};
    constant long strides_in1[] = {3,1};
    uint remaining = gid;
    uint idx_in0 = 0;
    uint idx_in1 = 0;
    for (int i = 1; i >= 0; i--) {
      uint coord = remaining % shape[i];
      idx_in0 += coord * strides_in0[i];
      idx_in1 += coord * strides_in1[i];
      remaining /= shape[i];
    }
    out[gid] = in0[idx_in0] * in1[idx_in1];
  }